Complete each local flag reset before entering a global barrier. This prevents a slower PE from overwriting the peer's first signal after the timed kernel launches begin.

Summary

• synchronize each PE after resetting its local ping-pong flag
• add a global barrier before starting each timed collective launch
• keep both operations outside the measured interval

Why

The timed loop resets each PE's symmetric flag immediately before a
collective launch. Unlike the warmup path and the other ping-pong tests in
this directory, it did not synchronize the reset across PEs.

This permits the following ordering:

1. PE 0 finishes its reset and launches the kernel.
2. PE 0 writes the first signal to PE 1's flag.
3. PE 1 performs its delayed reset and overwrites that signal with zero.
4. PE 1's kernel waits indefinitely for the lost signal.

A device-side atomic signal store, volatile store, or fence cannot repair
this race because the missing happens-before relationship is between the
host-initiated reset on one PE and the kernel launch on another PE.

The issue was reproduced on a downstream CUDA-compatible backend. This
change mirrors the reset ordering already used by the warmup path and the
other symmetric-flag ping-pong tests.

Testing

• downstream full library and perftest build
• 20 consecutive shmem_p_ping_pong_latency runs:
  • 10 with the legacy/static symmetric heap
  • 10 with the VMM symmetric heap
  • sizes 4 B through 16 KiB, 50 warmup iterations, 500 timed iterations
• three runs each of the four sibling ping-pong latency tests

Notes

The local device synchronization is needed to complete the local reset.
The following nvshmem_barrier_all() is needed to establish cross-PE
ordering. Neither operation is included in the latency measurement.